#include <stdio.h>
#include <algorithm>
#define THREADS_PER_BLOCK 16
#define DIVUP(m, n) ((m) / (n) + ((m) % (n) > 0))

// 打印数据的核函数
__global__ void printf_cuda(const float* const data_ptr,
    const int rows,
    const int cols){
    const int row = blockIdx.x * THREADS_PER_BLOCK + threadIdx.x;
    const int col = blockIdx.y * THREADS_PER_BLOCK + threadIdx.y;
    if(row>=rows||col>=cols)
        return;
    printf("data[%d][%d]=%f\n",row,col,*(data_ptr+row*cols+col));
}

// __device__: gpu函数调用gpu函数
__device__ float sigmoid(const float* const data_ptr){
    return 1.0/(1.0+exp(-*data_ptr));
}

__device__ float d_sigmoid(const float* const data_ptr){
    return exp(-*data_ptr)/((1.0+exp(-*data_ptr))*(1.0+exp(-*data_ptr)));
}

// __global__: cpu函数调用gpu函数
__global__ void sigmoid_forward(const float* const data_ptr,
    float* const output_ptr,
    const int rows,
    const int cols){
    const int row = blockIdx.x * THREADS_PER_BLOCK + threadIdx.x;
    const int col = blockIdx.y * THREADS_PER_BLOCK + threadIdx.y;
    if(row>=rows||col>=cols)
        return;
    *(output_ptr+row*cols+col) = sigmoid(data_ptr+row*cols+col);
}

__global__ void sigmoid_backward(const float* const data_ptr,
    float* const output_ptr,
    const int rows,
    const int cols){
    const int row = blockIdx.x * THREADS_PER_BLOCK + threadIdx.x;
    const int col = blockIdx.y * THREADS_PER_BLOCK + threadIdx.y;
    if(row>=rows||col>=cols)
        return;
    *(output_ptr+row*cols+col) = d_sigmoid(data_ptr+row*cols+col);
}

// 用于调用cuda核函数
void sigmoid_forward_launch(const float* const data_ptr,
    float* const output_ptr,
    const int rows,
    const int cols){
    dim3 blocks(DIVUP(rows,THREADS_PER_BLOCK),DIVUP(cols,THREADS_PER_BLOCK));
    dim3 threads(THREADS_PER_BLOCK,THREADS_PER_BLOCK);
    sigmoid_forward<<<blocks,threads>>>(data_ptr,output_ptr,rows,cols);
    cudaDeviceSynchronize();
//     printf_cuda<<<blocks,threads>>>(data_ptr,rows,cols);
//     printf_cuda<<<blocks,threads>>>(output_ptr,rows,cols);
}

void sigmoid_backward_launch(const float* const data_ptr,
    float* const output_ptr,
    const int rows,
    const int cols){
    dim3 blocks(DIVUP(rows,THREADS_PER_BLOCK),DIVUP(cols,THREADS_PER_BLOCK));
    dim3 threads(THREADS_PER_BLOCK,THREADS_PER_BLOCK);
    sigmoid_backward<<<blocks,threads>>>(data_ptr,output_ptr,rows,cols);
    cudaDeviceSynchronize();
}
